CUDA에서 카운트 변수 동기화 탐구

2025/01/16

오늘 아침 chatgpt와 질답 놀이를 했다. 기존에 알고 있던 내용들과 부합하지만 오류가 있을 수 있음.

GPU 공유메모리는 SRAM이다. 따라서 작음. 각각 다르지만 나처럼 대충 짜는 사람들은 50KB 안 넘긴다고 생각해야 됨.

SRAM이므로 전역램이랑 비교할 수 없음.

아래 코드를 참고. __syncthreads를 통해 공유 메모리 동기화 처리를 수행한다. 동기화는 공평해야 하므로 조건부 동기화가 일어나면 데드락에 빠질 수 있음을 주의.

atomicAdd는 모든 쓰레드가 경쟁함. __syncthreads()는 블록내 쓰레드끼리 경쟁함.

게다가 빠른 공유메모리를 사용하므로 결론은,

1. 공유메모리로 계산한 후 합산.
2. 블럭 수를 늘리면 경쟁이 줄어듬.
3. 블럭 수는 커널 실행시 정할 수 있음.
4. 블럭당 최대 스레드는 Coumpute Capability에 따라 다름. 현 시점에서 대충 1024라고 가정.
5. 블럭당 스레드 수가 적을 수록 유리함.
6. 블록당 스레드 수는 32의 배수가 유리함. 워프란 스레드 32개 묶음이고 32개 동시처리 SIMD 라고 생각하면 됨.
7. 최적 블럭 크기 계산은 cudaOccupancyMaxPotentialBlockSize
8. 블럭 수를 키우기면 하면 (블럭당 스레드 수를 줄이면) 최대 쿠다 코어수가 정해져 있으므로 대기 블록이 발생하고 병렬 처리 효율이 떨어짐
9. SM은 대충 100개이하, 스트림의 갯수는 내 맘대로지만 256개 이하라고 생각. 물론 코드로 알아내자.

참고: https://m.blog.naver.com/luku756/221600079153

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void count_character(const char *str, char target, int *count, int n) {
    // 스레드 ID 계산
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    // 카운트 저장용 공유 메모리
    __shared__ int blockCount[256];
    if (threadIdx.x == 0) blockCount[blockIdx.x] = 0;
    __syncthreads();

    // 유효한 인덱스에서 문자 검사
    if (idx < n && str[idx] == target) {
        atomicAdd(&blockCount[blockIdx.x], 1);
    }
    __syncthreads();

    // 결과를 전역 메모리에 저장
    if (threadIdx.x == 0) {
        atomicAdd(count, blockCount[blockIdx.x]);
    }
}

int main() {
    const char *input = "cuda is great and cuda is fast";
    char target = 'a';
    int n = strlen(input);

    // 문자열 및 카운트 변수 초기화
    char *d_str;
    int *d_count;
    int h_count = 0;

    cudaMalloc((void **)&d_str, n * sizeof(char));
    cudaMalloc((void **)&d_count, sizeof(int));

    cudaMemcpy(d_str, input, n * sizeof(char), cudaMemcpyHostToDevice);
    cudaMemcpy(d_count, &h_count, sizeof(int), cudaMemcpyHostToDevice);

    // 블록 및 스레드 구성
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    // 커널 실행
    count_character<<<blocksPerGrid, threadsPerBlock>>>(d_str, target, d_count, n);

    // 결과 복사 및 출력
    cudaMemcpy(&h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost);

    printf("The character '%c' appears %d times.\n", target, h_count);

    // 메모리 해제
    cudaFree(d_str);
    cudaFree(d_count);

    return 0;
}

#include <stdio.h>
#include <cuda_runtime.h>

int main() {
    int deviceCount;
    cudaDeviceProp deviceProp;

    // GPU 개수 확인
    cudaGetDeviceCount(&deviceCount);

    printf("Total devices: %d\n\n", deviceCount);

    for (int i = 0; i < deviceCount; ++i) {
        cudaGetDeviceProperties(&deviceProp, i);

        printf("Device %d: \n", i);
        printf("  SM count: %d\n", deviceProp.multiProcessorCount);
        printf("  Max threads per block: %d\n", deviceProp.maxThreadsPerBlock);
        printf("  Max block dimensions: %d, %d, %d\n", 
               deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1], deviceProp.maxThreadsDim[2]);
        printf("  Max grid dimensions: %d, %d, %d\n", 
               deviceProp.maxGridSize[0], deviceProp.maxGridSize[1], deviceProp.maxGridSize[2]);
        printf("\n");
    }

    return 0;
}

[t:/] is not "technology - root". dawnsea, rss